
#include "graph_gpu.h"

void InitGraphGpu(GraphGpu& dst, Graph& src) {
	size_t sz;
	

	sz = src.offs.size() * sizeof(int);
	CUDA_CHECK(cudaMalloc((void**)&dst.offs,  sz));
	CUDA_CHECK(cudaMemcpy(dst.offs, src.offs.data(), sz, cudaMemcpyHostToDevice));

	sz = src.adj.size() * sizeof(int);
	CUDA_CHECK(cudaMalloc((void**)&dst.adj, sz));
	CUDA_CHECK(cudaMemcpy(dst.adj, src.adj.data(), sz, cudaMemcpyHostToDevice));

	sz = src.probs.size() * sizeof(Float);
	CUDA_CHECK(cudaMalloc((void**)&dst.probs, sz));
	CUDA_CHECK(cudaMemcpy(dst.probs, src.probs.data(), sz, cudaMemcpyHostToDevice));

	sz = src.cores.size() * sizeof(int);
	CUDA_CHECK(cudaMalloc((void**)&dst.cores, sz));
	CUDA_CHECK(cudaMemcpy(dst.cores, src.cores.data(), sz, cudaMemcpyHostToDevice));

	dst.n = src.vertixNumber;
	dst.kmax = src.kmax;

}

void FreeGraphGpu(GraphGpu& graph) {
	CUDA_CHECK(cudaFree(graph.offs));
	CUDA_CHECK(cudaFree(graph.adj));
	CUDA_CHECK(cudaFree(graph.probs));
	CUDA_CHECK(cudaFree(graph.cores));
}

void InitGraphMetaGpu(GraphMetaGpu& dst, Graph& src) {
	size_t sz;

	//sz = sizeof(int);
	//CUDA_CHECK(cudaMalloc((void**)&dst.num, sz));
	//CUDA_CHECK(cudaMemset(dst.num, 0, sz));

	sz = src.vertixNumber * sizeof(int);

	CUDA_CHECK(cudaMalloc((void**)&dst.nodes, sz));
	CUDA_CHECK(cudaMemset(dst.nodes, 0, sz));
	CUDA_CHECK(cudaMalloc((void**)&dst.coreDegs, sz));
	CUDA_CHECK(cudaMemset(dst.coreDegs, 0, sz));
	CUDA_CHECK(cudaMalloc((void**)&dst.inQueue, sz));
	CUDA_CHECK(cudaMemset(dst.inQueue, 0, sz));

	sz = src.vertixNumber * sizeof(bool);
	CUDA_CHECK(cudaMalloc((void**)&dst.deleted, sz));
	CUDA_CHECK(cudaMemset(dst.deleted, 0, sz));
	CUDA_CHECK(cudaMalloc((void**)&dst.dirty, sz));
	CUDA_CHECK(cudaMemset(dst.dirty, 0, sz));

	sz = src.vertixNumber * sizeof(Float);

	CUDA_CHECK(cudaMalloc((void**)&dst.kprobs, sz));
	CUDA_CHECK(cudaMemset(dst.kprobs, 0, sz));
	CUDA_CHECK(cudaMalloc((void**)&dst.result, sz));
	CUDA_CHECK(cudaMemset(dst.result, 0, sz));
	CUDA_CHECK(cudaMalloc((void**)&dst.lb, sz));
	CUDA_CHECK(cudaMemset(dst.lb, 0, sz));

	sz = sizeof(Float);
	CUDA_CHECK(cudaMalloc((void**)&dst.minCurthres, sz));
	CUDA_CHECK(cudaMemset(dst.minCurthres, 0, sz));

	sz = sizeof(int);
	CUDA_CHECK(cudaMalloc((void**)&dst.remain, sz));
	CUDA_CHECK(cudaMemset(dst.remain, 0, sz));
	CUDA_CHECK(cudaMalloc((void**)&dst.flag, sz));
	CUDA_CHECK(cudaMemset(dst.flag, 0, sz));

	sz = sizeof(GlobalBuffer);
	CUDA_CHECK(cudaMalloc((void**)&dst.globalBuffer1, sz));
	CUDA_CHECK(cudaMemset(dst.globalBuffer1, 0, sz));
	CUDA_CHECK(cudaMalloc((void**)&dst.globalBuffer2, sz));
	CUDA_CHECK(cudaMemset(dst.globalBuffer2, 0, sz));
}

void FreeGraphMetaGpu(GraphMetaGpu& meta) {
	//CUDA_CHECK(cudaFree(meta.num));
	CUDA_CHECK(cudaFree(meta.nodes));
	CUDA_CHECK(cudaFree(meta.coreDegs));
	CUDA_CHECK(cudaFree(meta.deleted));
	CUDA_CHECK(cudaFree(meta.inQueue));
	CUDA_CHECK(cudaFree(meta.kprobs));
	CUDA_CHECK(cudaFree(meta.result));
	CUDA_CHECK(cudaFree(meta.lb));
	CUDA_CHECK(cudaFree(meta.dirty));
	CUDA_CHECK(cudaFree(meta.minCurthres));
	CUDA_CHECK(cudaFree(meta.remain));
	CUDA_CHECK(cudaFree(meta.flag));

	CUDA_CHECK(cudaFree(meta.globalBuffer1));
	CUDA_CHECK(cudaFree(meta.globalBuffer2));
}

void ClearGraphMetaGpu(GraphMetaGpu& dst, Graph& src) {
	size_t sz;
	sz = src.vertixNumber * sizeof(int);
	CUDA_CHECK(cudaMemset(dst.inQueue, 0, sz));

	sz = src.vertixNumber * sizeof(Float);
	CUDA_CHECK(cudaMemset(dst.kprobs, 0, sz));
	CUDA_CHECK(cudaMemset(dst.result, 0, sz));

	sz = src.vertixNumber * sizeof(bool);
	CUDA_CHECK(cudaMemset(dst.deleted, 0, sz));
	CUDA_CHECK(cudaMemset(dst.dirty, 0, sz));

	sz = sizeof(int);
	CUDA_CHECK(cudaMemset(dst.flag, 0, sz));
}

void SortByCore(GraphGpu& graph, GraphMetaGpu& meta) {
	int* idx = new int[graph.n];

	for (int i = 0; i < graph.n; ++i) idx[i] = i;

	CUDA_CHECK(cudaMemcpy(meta.nodes, idx, graph.n * sizeof(int), cudaMemcpyHostToDevice));

	thrust::device_ptr<int> idx_ptr(meta.nodes);
	thrust::device_vector<int> vecGpu(graph.cores, graph.cores + graph.n);
	thrust::sort_by_key(vecGpu.begin(), vecGpu.end(), idx_ptr, thrust::greater<int>());

	delete[] idx;
}

__global__ void ComputeCoreDeg(int k, int num,
	int* __restrict__ adj,
	int* __restrict__ offs,
	int* __restrict__ cores,
	int* __restrict__ nodes,
	int* __restrict__ degs) {

	int globalThreadNum = gridDim.x * blockDim.x;
	int globalThreadId = blockDim.x * blockIdx.x + threadIdx.x;

	//if (globalThreadId == 0) {
	//	for (int i = 0; i < 9; i++) {
	//		CUDA_INFO("cores[%d]=%d", i, cores[i]);
	//	}
	//}
	int u, v, left, right;
	for (int i = globalThreadId; i < num; i += globalThreadNum) {
		u = nodes[i];
		left = offs[u];
		right = offs[u + 1];
		for (int j = left; j < right; ++j) {
			v = adj[j];
			if (cores[v] < k) {
				degs[u] = j - left;
				//CUDA_INFO("k=%d,degs[%d]=%d j=%d left=%d v=%d %d", k, u, degs[u], j, left, v, cores[v]);
				break;
			}
			if (j == right - 1) {
				degs[u] = right - left;
			}
		}
	}
}